package avl.sv.shared.model.featureGenerator.jocl;

import java.awt.image.BufferedImage;
import java.awt.image.DataBufferByte;
import java.io.BufferedReader;
import java.io.IOException;
import java.io.InputStream;
import java.io.InputStreamReader;
import java.util.ArrayList;
import org.jocl.*;
import org.jocl.Pointer;
import org.jocl.Sizeof;
import org.jocl.cl_command_queue;
import org.jocl.cl_context;
import org.jocl.cl_context_properties;
import org.jocl.cl_device_id;
import org.jocl.cl_event;
import org.jocl.cl_kernel;
import org.jocl.cl_mem;
import org.jocl.cl_platform_id;
import org.jocl.cl_program;

public class FeatureGeneratorJOCL {
    
    private cl_command_queue commandQueue;
    private cl_context context = null;
    private cl_kernel   clMeansKernel, 
                        clContrastKernel, 
                        clMomentsKernel, 
                        clRGB_DeInterlaceKernel, 
                        clCorrelationKernel, 
                        clHomogeneityKernel, 
                        clRGB_ToGrayKernel, 
                        clKurtosisSkewnessKernal,
                        clMagnitudeKernel,
                        clSumKernel, 
                        clRGB_DeInterlaceToHSVKernel;
    
    ArrayList<Plane> planes = new ArrayList<>();
    private long deviceMemoryTotal = -1;

    private final int platformIndex, deviceIndex;
    
    public FeatureGeneratorJOCL(final int platformIndex, final int deviceIndex) {
        this.platformIndex = platformIndex;
        this.deviceIndex = deviceIndex;
        // Setup Default features
        for (Plane.Names p:new Plane.Names[]{Plane.Names.Red, Plane.Names.Green, Plane.Names.Blue}){
            ArrayList<Feature> features = new ArrayList<>();
            features.add(new Feature(Feature.Names.Mean));
            planes.add(new Plane(p, features));           
        }
    }
    
    private void setupJOCL() {
        // The platform, device type and device number
        // that will be used
        final long deviceType = CL.CL_DEVICE_TYPE_ALL;

        // Enable exceptions and subsequently omit error checks in this sample
        CL.setExceptionsEnabled(true);

        // Obtain the number of platforms
        int numPlatformsArray[] = new int[1];
        CL.clGetPlatformIDs(0, null, numPlatformsArray);
        int numPlatforms = numPlatformsArray[0];

        // Obtain a platform ID
        cl_platform_id platforms[] = new cl_platform_id[numPlatforms];
        CL.clGetPlatformIDs(platforms.length, platforms, null);
        cl_platform_id platform = platforms[platformIndex];

        // Initialize the context properties
        cl_context_properties contextProperties = new cl_context_properties();
        contextProperties.addProperty(CL.CL_CONTEXT_PLATFORM, platform);

        // Obtain the number of devices for the platform
        int numDevicesArray[] = new int[1];
        CL.clGetDeviceIDs(platform, deviceType, 0, null, numDevicesArray);
        int numDevices = numDevicesArray[0];

        // Obtain a device ID 
        cl_device_id devices[] = new cl_device_id[numDevices];
        CL.clGetDeviceIDs(platform, deviceType, numDevices, devices, null);
        cl_device_id device = devices[deviceIndex];
        
        // Create a context for the selected device
        context = CL.clCreateContext(contextProperties, 1, new cl_device_id[]{device}, null, null, null);

        // Create a command-queue for the selected device
        commandQueue = CL.clCreateCommandQueue(context, device, 0, null);
        
        clMeansKernel = loadAndCompileKernel("Means_kernal.cl", "means_kernal");
        clContrastKernel = loadAndCompileKernel("Contrast_kernal.cl", "contrast_kernal");
        clMomentsKernel = loadAndCompileKernel("Moments_kernal.cl", "moments_kernal");
        clRGB_DeInterlaceKernel = loadAndCompileKernel("RGB_DeInterlace_kernal.cl", "rgb_DeInterlace_kernal");
        clRGB_DeInterlaceToHSVKernel = loadAndCompileKernel("RGB_DeInterlaceToHSV_kernal.cl", "rgb_DeInterlaceToHSV_kernal");
        clCorrelationKernel = loadAndCompileKernel("Correlation_kernal.cl", "correlation_kernal");
        clHomogeneityKernel = loadAndCompileKernel("Homogeneity_kernal.cl", "homogeneity_kernal");
        clRGB_ToGrayKernel = loadAndCompileKernel("rgb_ToGray_kernal.cl", "rgb_ToGray_kernal");
        clKurtosisSkewnessKernal = loadAndCompileKernel("KurtosisSkewness_kernal.cl", "kurtosisSkewness_kernal");
        clMagnitudeKernel = loadAndCompileKernel("Magnitude_kernal.cl", "magnitude_kernal");
        clSumKernel = loadAndCompileKernel("Sum_kernal.cl", "sum_kernal");
        
    }
    
    private class CL_Feature{
        Feature.Names name;
        cl_mem cl_data;
        cl_event cl_event = new cl_event();
        ArrayList<cl_event> waitForEvents = new ArrayList<>();
    }
    
    private class CL_RGB_Set {

        final int w, h, numOfImages;
        final int bands = 3;
        final cl_mem imgMem;

        public CL_RGB_Set(int numOfImages, int width, int height, byte data[], ArrayList<cl_event> waitForEvents) {
            w = width;
            h = height;
            this.numOfImages = numOfImages;
            imgMem = CL.clCreateBuffer(context, CL.CL_MEM_READ_ONLY, w * h * numOfImages * bands, null, null);
            cl_event event = new cl_event();
            CL.clEnqueueWriteBuffer(commandQueue, imgMem, false, 0, data.length, Pointer.to(data), 0, null, event);
            waitForEvents.add(event);
        }

        public void free() {
            CL.clReleaseMemObject(imgMem);
        }
    }

    private class CL_PlaneSet {

        Plane.Names planeName;
        final int w, h, numOfImages;
        final cl_mem imgMem;

        public CL_PlaneSet(Plane.Names name, int w, int h, int numOfImages, cl_mem imgMem) {
            this.planeName = name;
            this.w = w;
            this.h = h;
            this.numOfImages = numOfImages;
            this.imgMem = imgMem;
        }

        public void free() {
            CL.clReleaseMemObject(imgMem);
        }
    }
    
    private class CL_FeatureSet {

        final ArrayList<CL_Feature> cl_features = new ArrayList<>();
        final CL_PlaneSet cL_Plane;
        private ArrayList<cl_event> waitForEvents;

        public CL_FeatureSet(CL_PlaneSet cL_Plane, ArrayList<cl_event> waitForEvents) {
            this.cL_Plane = cL_Plane;
            this.waitForEvents = (ArrayList<cl_event>) waitForEvents.clone();
        }

        public ArrayList<CL_Feature> getFeatures(ArrayList<Feature.Names> featureNames) {
            ArrayList<CL_Feature> outFeatures = new ArrayList<>();
            for (Feature.Names featureName : featureNames) {
                CL_Feature feature = getFeature(featureName, waitForEvents);
                outFeatures.add(feature);
            }
            return outFeatures;
        }

        public void free() {
            for (CL_Feature cl_feature : cl_features) {
                CL.clReleaseMemObject(cl_feature.cl_data);
            }
        }

        public CL_Feature getFeature(Feature.Names featureName, ArrayList<cl_event> waitForEvents) {
            for (CL_Feature feature : cl_features) {
                if (feature.name.equals(featureName)) {
                    return feature;
                }
            }
            ArrayList<CL_Feature> tempFeatures;
            CL_Feature tempFeature0, tempFeature1, tempFeature2;
            switch (featureName) {
                case Mean:
                    CL_Feature feature = mean(cL_Plane, waitForEvents);
                    cl_features.add(feature);
                    break;
                case Moment2: // Fall through
                case Moment3:
                    tempFeature0 = getFeature(Feature.Names.Mean, waitForEvents);
                    tempFeatures = moments(cL_Plane, waitForEvents, tempFeature0);
                    cl_features.addAll(tempFeatures);
                    break;
                case Kurtosis: // Fall through
                case Skewness:
                    tempFeature0 = getFeature(Feature.Names.Mean, waitForEvents);
                    tempFeatures = kurtosisSkewness(cL_Plane, waitForEvents, tempFeature0);
                    cl_features.addAll(tempFeatures);
                    break;
                case ContrastD: // Fall through
                case ContrastR:
                    
                    tempFeatures = contrast(cL_Plane, waitForEvents);
                    cl_features.addAll(tempFeatures);
                    break;
                case HomogeneityD: // Fall through
                case HomogeneityR:
                    tempFeatures = homogeneity(cL_Plane, waitForEvents);
                    cl_features.addAll(tempFeatures);
                    break;
                case CorrelationD: // Fall through
                case CorrelationR:
                    tempFeature0 = getFeature(Feature.Names.Mean, waitForEvents);
                    tempFeature1 = getFeature(Feature.Names.Moment2, waitForEvents);
                    tempFeatures = correlation(cL_Plane, waitForEvents, tempFeature0, tempFeature1);
                    cl_features.addAll(tempFeatures);
                    break;
                case CorrelationMagnitude:
                    tempFeature0 = getFeature(Feature.Names.CorrelationD, waitForEvents);
                    tempFeature1 = getFeature(Feature.Names.CorrelationR, waitForEvents);
                    tempFeature2 = magnitude(cL_Plane, waitForEvents, tempFeature0, tempFeature1);
                    tempFeature2.name = Feature.Names.CorrelationMagnitude;
                    cl_features.add(tempFeature2);
                    break;
                case HomogeneityMagnitude:
                    tempFeature0 = getFeature(Feature.Names.HomogeneityD, waitForEvents);
                    tempFeature1 = getFeature(Feature.Names.HomogeneityR, waitForEvents);
                    tempFeature2 = magnitude(cL_Plane, waitForEvents, tempFeature0, tempFeature1);
                    tempFeature2.name = Feature.Names.HomogeneityMagnitude;
                    cl_features.add(tempFeature2);
                    break;
                case ContrastMagnitude:
                    tempFeature0 = getFeature(Feature.Names.ContrastD, waitForEvents);
                    tempFeature1 = getFeature(Feature.Names.ContrastR, waitForEvents);
                    tempFeature2 = magnitude(cL_Plane, waitForEvents, tempFeature0, tempFeature1);
                    tempFeature2.name = Feature.Names.ContrastMagnitude;
                    cl_features.add(tempFeature2);
                    break;
                case CorrelationSum:
                    tempFeature0 = getFeature(Feature.Names.CorrelationD, waitForEvents);
                    tempFeature1 = getFeature(Feature.Names.CorrelationR, waitForEvents);
                    tempFeature2 = fSum(cL_Plane, waitForEvents, tempFeature0, tempFeature1);
                    tempFeature2.name = Feature.Names.CorrelationSum;
                    cl_features.add(tempFeature2);
                    break;
                case HomogeneitySum:
                    tempFeature0 = getFeature(Feature.Names.HomogeneityD, waitForEvents);
                    tempFeature1 = getFeature(Feature.Names.HomogeneityR, waitForEvents);
                    tempFeature2 = fSum(cL_Plane, waitForEvents, tempFeature0, tempFeature1);
                    tempFeature2.name = Feature.Names.HomogeneitySum;
                    cl_features.add(tempFeature2);
                    break;
                case ContrastSum:
                    tempFeature0 = getFeature(Feature.Names.ContrastD, waitForEvents);
                    tempFeature1 = getFeature(Feature.Names.ContrastR, waitForEvents);
                    tempFeature2 = fSum(cL_Plane, waitForEvents, tempFeature0, tempFeature1);
                    tempFeature2.name = Feature.Names.ContrastSum;
                    cl_features.add(tempFeature2);
                    break;
                default:
                    throw new IllegalArgumentException(featureName.name() + " is Not a valid feature name");
            }
            for (CL_Feature feature : cl_features) {
                if (feature.name.equals(featureName)) {
                    return feature;
                }
            }
            throw new Error("Failed to generate a feature");
        }
    }

    private ArrayList<CL_PlaneSet> getRequiredCL_PlaneSets(BufferedImage[] imgs, ArrayList<cl_event> waitForEvents){
        final CL_RGB_Set imgSetRGB = imagesToGPU(imgs, waitForEvents);
        ArrayList<CL_PlaneSet> cl_planeSetsTemp = new ArrayList<>();
        ArrayList<Plane.Names> planesToGenerate = new ArrayList<>();
        for (Plane plane:planes){
            planesToGenerate.add(plane.planeName);
        }
        while (!planesToGenerate.isEmpty()){
            if (planesToGenerate.contains(Plane.Names.Red) | planesToGenerate.contains(Plane.Names.Green) | planesToGenerate.contains(Plane.Names.Blue)){
                cl_planeSetsTemp.addAll(rgbDeInterlace(imgSetRGB, waitForEvents));
                planesToGenerate.remove(Plane.Names.Red);
                planesToGenerate.remove(Plane.Names.Green);
                planesToGenerate.remove(Plane.Names.Blue);
            }
            if (planesToGenerate.contains(Plane.Names.Hue) | planesToGenerate.contains(Plane.Names.Saturation) | planesToGenerate.contains(Plane.Names.Value)){
                cl_planeSetsTemp.addAll(rgbDeInterlaceToHSV(imgSetRGB, waitForEvents));
                planesToGenerate.remove(Plane.Names.Hue);
                planesToGenerate.remove(Plane.Names.Saturation);
                planesToGenerate.remove(Plane.Names.Value);
            }
            if (planesToGenerate.contains(Plane.Names.Gray)){
                cl_planeSetsTemp.addAll(rgbToGray(imgSetRGB, waitForEvents));
                planesToGenerate.remove(Plane.Names.Gray);
            }
        }
        
        ArrayList<CL_PlaneSet> cl_planeSets = new ArrayList<>();
        for (Plane plane:planes){
            for (int i = 0; i < cl_planeSetsTemp.size(); i++) {
                if (plane.planeName == cl_planeSetsTemp.get(i).planeName){
                    CL_PlaneSet temp = cl_planeSetsTemp.get(i);
                    cl_planeSets.add(temp);
                    break;
                }
            }
        }
        cl_planeSetsTemp.removeAll(cl_planeSets);
        for (CL_PlaneSet planeSet:cl_planeSetsTemp){
            planeSet.free();
        }
        imgSetRGB.free();
        return cl_planeSets;
    }
    
    public double[][] getFeaturesForImages(BufferedImage[] imgs) {
        if (context == null){
            setupJOCL();
        }
        CL.setExceptionsEnabled(true);
        
        double featuresData[][] = new double[imgs.length][getNumberOfFeatures()];
        
        ArrayList<cl_event> waitForEvents = new ArrayList<>();
        final ArrayList<CL_PlaneSet> cl_planeSets = getRequiredCL_PlaneSets(imgs, waitForEvents);
        int featureIdx = 0;
        for (final CL_PlaneSet cl_planeSet : cl_planeSets) {
            // print_cl_mem_byte(plane.imgMem, plane.w*plane.h, waitForEvents);
            ArrayList<Feature.Names> featureNames = new ArrayList<>();
            for (Plane plane:planes){
                if (plane.planeName.equals(cl_planeSet.planeName)){
                    ArrayList<Feature> temp = plane.features;
                    for (Feature t:temp){
                        featureNames.add(t.featureName);
                    }
                }
            }
            ArrayList<Feature> features = generateFeatures(cl_planeSet, featureNames, waitForEvents);
            for (int f = 0; f < featureNames.size(); f++){
                Feature.Names featureName = featureNames.get(f);
                for (Feature feature:features){
                    if (feature.featureName.equals(featureName)){
                        for (int s = 0; s < imgs.length; s++){
                            featuresData[s][featureIdx] = feature.data[s];
                        }
                        featureIdx++;
                    }
                }
            }
            cl_planeSet.free();
        }    
        return featuresData;
    }

    public String[] getFeatureNames() {
        ArrayList<String> names = new ArrayList<>();
        for (Plane plane:planes){
            for (Feature feature:plane.features){
                names.add(plane.planeName + "->" + feature.featureName);
            }
        }
        return names.toArray(new String[names.size()]);
    }
    
    public void setFeatureNames(String[] featureNames) {
        planes = new ArrayList<>();
        for (String featureString:featureNames){
            
            String[] temp = featureString.split("->");
            if (temp.length != 2){
                continue;
            }
            Plane.Names planeName = Plane.Names.valueOf(temp[0]);
            Feature.Names featureName = Feature.Names.valueOf(temp[1]);
            if ((planeName != null) && (featureName != null)){
                boolean featureAdded = false;
                addFeature:
                for (Plane plane:planes){
                    if (plane.planeName.equals(planeName)){
                        for (Feature feature : plane.features) {
                            if (feature.featureName.equals(featureName)) {
                                featureAdded = true;
                                break addFeature;
                            }
                        }
                        plane.features.add(new Feature(featureName));
                        featureAdded = true;
                        break addFeature;
                    }
                }
                if (featureAdded){
                    continue;
                }
                ArrayList<Feature> features = new ArrayList<>();
                features.add(new Feature(featureName));
                planes.add(new Plane(planeName, features));
            }
        }
    }

    public void setPlanes(ArrayList<Plane> planes){
        this.planes = planes;
    }
    
    public int getNumberOfFeatures() {
        int count = 0; 
        for (Plane plane:planes){
            count += plane.features.size();
        }
        return count;
    }

    private ArrayList<Feature> generateFeatures(CL_PlaneSet imgPlane, ArrayList<Feature.Names> featureNames, ArrayList<cl_event> waitForEvents) {

        CL_FeatureSet featureSet = new CL_FeatureSet(imgPlane, waitForEvents);
                
        ArrayList<CL_Feature> CL_features = featureSet.getFeatures(featureNames);

        // Copy out features
        for (CL_Feature feature : CL_features) {
            waitForEvents.add(feature.cl_event);
        }
        cl_event[] waitForEventsArray = waitForEvents.toArray(new cl_event[waitForEvents.size()]);
        ArrayList<Feature> features = new ArrayList<>();
        for (int i = 0; i < CL_features.size(); i++) {
            CL_Feature cl_feature = CL_features.get(i);
            float[] featureData = new float[imgPlane.numOfImages];
            cl_event readFeatures = new cl_event();
            CL.clEnqueueReadBuffer(commandQueue, cl_feature.cl_data, 
                    CL.CL_TRUE, 0, imgPlane.numOfImages * Sizeof.cl_float,
                    Pointer.to(featureData), waitForEventsArray.length, waitForEventsArray, readFeatures);
            waitForEvents.add(readFeatures);
            Feature feature = new Feature(cl_feature.name);
            feature.data = featureData;
            features.add(feature);
        }
        waitForEventsArray = waitForEvents.toArray(new cl_event[waitForEvents.size()]);
        CL.clWaitForEvents(waitForEventsArray.length, waitForEventsArray);
        featureSet.free();
        return features;
    }

    private CL_RGB_Set imagesToGPU(final BufferedImage imgs[], ArrayList<cl_event> waitForEvents) {
        final int w = imgs[0].getWidth();
        final int h = imgs[0].getHeight();
        final int bands = 3;
        final byte b[] = new byte[w * h * imgs.length * bands];
        for (int i = 0; i < imgs.length; i++) {
            final int offset = i * w * h * bands;
            final byte data[] = ((DataBufferByte) imgs[i].getRaster().getDataBuffer()).getData();
            System.arraycopy(data, 0, b, offset, w * h * bands);
        }
        CL_RGB_Set imgSet = new CL_RGB_Set(imgs.length, imgs[0].getWidth(), imgs[0].getHeight(), b, waitForEvents);
        return imgSet;
    }

    private static long round(long groupSize, long globalSize) {
        long r = globalSize % groupSize;
        if (r == 0) {
            return globalSize;
        } else {
            return globalSize + groupSize - r;
        }
    }

    private static String readKernel(InputStream kernalStream) {
        try {
            BufferedReader br = new BufferedReader(new InputStreamReader(kernalStream));
            StringBuilder sb = new StringBuilder();
            String line;
            while (true) {
                line = br.readLine();
                if (line == null) {
                    break;
                }
                sb.append(line).append("\n");
            }
            return sb.toString();
        } catch (IOException e) {
            System.exit(1);
            return null;
        }
    }

    private cl_kernel loadAndCompileKernel(String kernelFileName, String kernelName) {
        // Create the OpenCL kernel from the program
        InputStream kernelStream = getClass().getResourceAsStream(kernelFileName);
        String source = readKernel(kernelStream);
        cl_program program = CL.clCreateProgramWithSource(context, 1,
                new String[]{source}, null, null);
        String compileOptions = "-cl-mad-enable";
        CL.clBuildProgram(program, 0, null, compileOptions, null, null);
        cl_kernel out = CL.clCreateKernel(program, kernelName, null);
        CL.clReleaseProgram(program);
        return out;
    }

    @Override
    protected void finalize() throws Throwable {       
        CL.clReleaseKernel(clMeansKernel);
        CL.clReleaseKernel(clContrastKernel);
        CL.clReleaseKernel(clMomentsKernel);
        CL.clReleaseKernel(clRGB_DeInterlaceKernel);
        CL.clReleaseKernel(clCorrelationKernel);
        CL.clReleaseKernel(clHomogeneityKernel);
        CL.clReleaseKernel(clRGB_ToGrayKernel);        
        CL.clReleaseCommandQueue(commandQueue);
        CL.clReleaseContext(context);
        super.finalize();
    }
       
    private ArrayList<CL_PlaneSet> rgbDeInterlace(CL_RGB_Set imgs, ArrayList<cl_event> waitForEvents) {
        cl_kernel kernel = clRGB_DeInterlaceKernel;
        int numelPixels = imgs.w * imgs.h;
        int params[] = {imgs.numOfImages, numelPixels, 0, 0};
        cl_mem outR = CL.clCreateBuffer(context, CL.CL_MEM_READ_WRITE, imgs.numOfImages * numelPixels, null, null);
        cl_mem outG = CL.clCreateBuffer(context, CL.CL_MEM_READ_WRITE, imgs.numOfImages * numelPixels, null, null);
        cl_mem outB = CL.clCreateBuffer(context, CL.CL_MEM_READ_WRITE, imgs.numOfImages * numelPixels, null, null);
        CL.clSetKernelArg(kernel, 0, Sizeof.cl_mem, Pointer.to(imgs.imgMem));
        CL.clSetKernelArg(kernel, 1, Sizeof.cl_mem, Pointer.to(outR));
        CL.clSetKernelArg(kernel, 2, Sizeof.cl_mem, Pointer.to(outG));
        CL.clSetKernelArg(kernel, 3, Sizeof.cl_mem, Pointer.to(outB));
        CL.clSetKernelArg(kernel, 4, Sizeof.cl_int2, Pointer.to(params));

        cl_event event = new cl_event();
        long localWorkSize[] = {imgs.w};
        long globalWorkSize[] = {localWorkSize[0] * imgs.numOfImages};
        cl_event[] waitForEventsArray = waitForEvents.toArray(new cl_event[waitForEvents.size()]);
        CL.clEnqueueNDRangeKernel(commandQueue, kernel, 1, null, globalWorkSize, localWorkSize, waitForEventsArray.length, waitForEventsArray, event);
        waitForEvents.add(event);

        ArrayList<CL_PlaneSet> rgbPlanes = new ArrayList<>();
        rgbPlanes.add(new CL_PlaneSet(Plane.Names.Red, imgs.w, imgs.h, imgs.numOfImages, outR));
        rgbPlanes.add(new CL_PlaneSet(Plane.Names.Green, imgs.w, imgs.h, imgs.numOfImages, outG));
        rgbPlanes.add(new CL_PlaneSet(Plane.Names.Blue, imgs.w, imgs.h, imgs.numOfImages, outB));
        return rgbPlanes;
    }

    private ArrayList<CL_PlaneSet> rgbDeInterlaceToHSV(CL_RGB_Set imgs, ArrayList<cl_event> waitForEvents) {
        cl_kernel kernel = clRGB_DeInterlaceToHSVKernel;
        int numelPixels = imgs.w * imgs.h;
        int params[] = {imgs.numOfImages, numelPixels, 0, 0};
        cl_mem outH = CL.clCreateBuffer(context, CL.CL_MEM_READ_WRITE, imgs.numOfImages * numelPixels, null, null);
        cl_mem outS = CL.clCreateBuffer(context, CL.CL_MEM_READ_WRITE, imgs.numOfImages * numelPixels, null, null);
        cl_mem outV = CL.clCreateBuffer(context, CL.CL_MEM_READ_WRITE, imgs.numOfImages * numelPixels, null, null);
        CL.clSetKernelArg(kernel, 0, Sizeof.cl_mem, Pointer.to(imgs.imgMem));
        CL.clSetKernelArg(kernel, 1, Sizeof.cl_mem, Pointer.to(outH));
        CL.clSetKernelArg(kernel, 2, Sizeof.cl_mem, Pointer.to(outS));
        CL.clSetKernelArg(kernel, 3, Sizeof.cl_mem, Pointer.to(outV));
        CL.clSetKernelArg(kernel, 4, Sizeof.cl_int2, Pointer.to(params));

        cl_event event = new cl_event();
        long localWorkSize[] = {imgs.w};
        long globalWorkSize[] = {localWorkSize[0] * imgs.numOfImages};
        cl_event[] waitForEventsArray = waitForEvents.toArray(new cl_event[waitForEvents.size()]);
        CL.clEnqueueNDRangeKernel(commandQueue, kernel, 1, null, globalWorkSize, localWorkSize, waitForEventsArray.length, waitForEventsArray, event);
        waitForEvents.add(event);

        ArrayList<CL_PlaneSet> rgbPlanes = new ArrayList<>();
        rgbPlanes.add(new CL_PlaneSet(Plane.Names.Hue, imgs.w, imgs.h, imgs.numOfImages, outH));
        rgbPlanes.add(new CL_PlaneSet(Plane.Names.Saturation, imgs.w, imgs.h, imgs.numOfImages, outS));
        rgbPlanes.add(new CL_PlaneSet(Plane.Names.Value, imgs.w, imgs.h, imgs.numOfImages, outV));
        return rgbPlanes;
    }
    
    private ArrayList<CL_PlaneSet> rgbToGray(CL_RGB_Set imgs, ArrayList<cl_event> waitForEvents) {
        cl_kernel kernel = clRGB_ToGrayKernel;
        int numelPixels = imgs.w * imgs.h;
        int params[] = {imgs.numOfImages, numelPixels, 0, 0};
        cl_mem grays = CL.clCreateBuffer(context, CL.CL_MEM_READ_WRITE, imgs.numOfImages * numelPixels, null, null);
        CL.clSetKernelArg(kernel, 0, Sizeof.cl_mem, Pointer.to(imgs.imgMem));
        CL.clSetKernelArg(kernel, 1, Sizeof.cl_mem, Pointer.to(grays));
        CL.clSetKernelArg(kernel, 2, Sizeof.cl_int2, Pointer.to(params));

        cl_event event = new cl_event();
        long localWorkSize[] = {imgs.w};
        long globalWorkSize[] = {localWorkSize[0] * imgs.numOfImages};
        cl_event[] waitForEventsArray = waitForEvents.toArray(new cl_event[waitForEvents.size()]);
        CL.clEnqueueNDRangeKernel(commandQueue, kernel, 1, null, globalWorkSize, localWorkSize, waitForEventsArray.length, waitForEventsArray, event);
        waitForEvents.add(event);

        ArrayList<CL_PlaneSet> rgbPlanes = new ArrayList<>();
        rgbPlanes.add(new CL_PlaneSet(Plane.Names.Gray, imgs.w, imgs.h, imgs.numOfImages, grays));
        return rgbPlanes;
    }
    
    private CL_Feature mean(final CL_PlaneSet imgs, ArrayList<cl_event> waitForEvents) {
        cl_kernel kernel = clMeansKernel;
        CL_Feature feature = new CL_Feature();
        feature.name = Feature.Names.Mean;

        int numelPixels = imgs.w * imgs.h;
        int params[] = {imgs.numOfImages, numelPixels};
        feature.cl_data = CL.clCreateBuffer(context, CL.CL_MEM_READ_WRITE, imgs.numOfImages * Sizeof.cl_float, null, null);
        CL.clSetKernelArg(kernel, 0, Sizeof.cl_mem, Pointer.to(imgs.imgMem));
        CL.clSetKernelArg(kernel, 1, Sizeof.cl_mem, Pointer.to(feature.cl_data));
        CL.clSetKernelArg(kernel, 2, Sizeof.cl_int2, Pointer.to(params));

        long localWorkSize[] = {1};
        long globalWorkSize[] = {imgs.numOfImages};
        cl_event[] waitForEventsArray = waitForEvents.toArray(new cl_event[waitForEvents.size()]);
        CL.clEnqueueNDRangeKernel(commandQueue, kernel, 1, null, globalWorkSize, localWorkSize, waitForEventsArray.length, waitForEventsArray, feature.cl_event);
        waitForEvents.add(feature.cl_event);
//        JOCL_Printer.print_cl_mem_Float(commandQueue, feature.cl_data, imgs.numOfImages, waitForEvents);
//        JOCL_Printer.print_cl_mem_byte(commandQueue, imgs.imgMem, imgs.h*imgs.w*imgs.numOfImages, waitForEvents);
        return feature;
    }

    private ArrayList<CL_Feature> moments(CL_PlaneSet imgs, ArrayList<cl_event> waitForEvents, CL_Feature means) {
        cl_kernel kernel = clMomentsKernel;
        CL_Feature moment2 = new CL_Feature();
        CL_Feature moment3 = new CL_Feature();

        moment2.name = Feature.Names.Moment2;
        moment3.name = Feature.Names.Moment3;

        int numelPixels = imgs.w * imgs.h;
        int params[] = {imgs.numOfImages, numelPixels, imgs.w, 0};

        moment2.cl_data = CL.clCreateBuffer(context, CL.CL_MEM_READ_WRITE, imgs.numOfImages * Sizeof.cl_float, null, null);
        moment3.cl_data = CL.clCreateBuffer(context, CL.CL_MEM_READ_WRITE, imgs.numOfImages * Sizeof.cl_float, null, null);

        CL.clSetKernelArg(kernel, 0, Sizeof.cl_mem, Pointer.to(imgs.imgMem));
        CL.clSetKernelArg(kernel, 1, Sizeof.cl_mem, Pointer.to(means.cl_data));
        CL.clSetKernelArg(kernel, 2, Sizeof.cl_mem, Pointer.to(moment2.cl_data));
        CL.clSetKernelArg(kernel, 3, Sizeof.cl_mem, Pointer.to(moment3.cl_data));
        CL.clSetKernelArg(kernel, 4, Sizeof.cl_int4, Pointer.to(params));

        cl_event event = new cl_event();
        long localWorkSize[] = {imgs.w, 1};
        long globalWorkSize[] = {localWorkSize[0] * imgs.numOfImages, 1};
        cl_event[] waitForEventsArray = waitForEvents.toArray(new cl_event[waitForEvents.size()]);
        CL.clEnqueueNDRangeKernel(commandQueue, kernel, 1, null, globalWorkSize, localWorkSize, waitForEventsArray.length, waitForEventsArray, event);
        moment2.cl_event = event;
        moment3.cl_event = event;

        ArrayList<CL_Feature> features = new ArrayList<>();
        features.add(moment2);
        features.add(moment3);

        return features;
    }

    private ArrayList<CL_Feature> contrast(CL_PlaneSet imgs, ArrayList<cl_event> waitForEvents) {
        cl_kernel kernel = clContrastKernel;
        CL_Feature outR = new CL_Feature();
        CL_Feature outD = new CL_Feature();

        outR.name = Feature.Names.ContrastR;
        outD.name = Feature.Names.ContrastD;

        outR.cl_data = CL.clCreateBuffer(context, CL.CL_MEM_READ_WRITE, imgs.numOfImages * Sizeof.cl_float, null, null);
        outD.cl_data = CL.clCreateBuffer(context, CL.CL_MEM_READ_WRITE, imgs.numOfImages * Sizeof.cl_float, null, null);

        int numelPixels = imgs.w * imgs.h;
        int params[] = {imgs.numOfImages, numelPixels, imgs.w, 0};
        CL.clSetKernelArg(kernel, 0, Sizeof.cl_mem, Pointer.to(imgs.imgMem));
        CL.clSetKernelArg(kernel, 1, Sizeof.cl_mem, Pointer.to(outR.cl_data));
        CL.clSetKernelArg(kernel, 2, Sizeof.cl_mem, Pointer.to(outD.cl_data));
        CL.clSetKernelArg(kernel, 3, Sizeof.cl_float * imgs.w, null);
        CL.clSetKernelArg(kernel, 4, Sizeof.cl_float * imgs.w, null);
        CL.clSetKernelArg(kernel, 5, Sizeof.cl_int * params.length, Pointer.to(params));

        cl_event event = new cl_event();

        long localWorkSize[] = {imgs.w};
        long globalWorkSize[] = {localWorkSize[0] * imgs.numOfImages};
        cl_event[] waitForEventsArray = waitForEvents.toArray(new cl_event[waitForEvents.size()]);
        CL.clEnqueueNDRangeKernel(commandQueue, kernel, 1, null, globalWorkSize, localWorkSize, waitForEventsArray.length, waitForEventsArray, event);
        outR.cl_event = event;
        outD.cl_event = event;

//        waitForEvents.add(event);
//        print_cl_mem_Float(outR.cl_data, imgs.numOfImages, waitForEvents);
//        print_cl_mem_byte(imgs.imgMem, 10, waitForEvents);
        ArrayList<CL_Feature> features = new ArrayList<>();
        features.add(outR);
        features.add(outD);
        return features;
    }

    private ArrayList<CL_Feature> homogeneity(CL_PlaneSet imgs, ArrayList<cl_event> waitForEvents) {
        cl_kernel kernel = clHomogeneityKernel;
        CL_Feature outR = new CL_Feature();
        CL_Feature outD = new CL_Feature();

        outR.name = Feature.Names.HomogeneityR;
        outD.name = Feature.Names.HomogeneityD;

        outR.cl_data = CL.clCreateBuffer(context, CL.CL_MEM_READ_WRITE, imgs.numOfImages * Sizeof.cl_float, null, null);
        outD.cl_data = CL.clCreateBuffer(context, CL.CL_MEM_READ_WRITE, imgs.numOfImages * Sizeof.cl_float, null, null);

        int numelPixels = imgs.w * imgs.h;
        int params[] = {imgs.numOfImages, numelPixels, imgs.w, 0};
        CL.clSetKernelArg(kernel, 0, Sizeof.cl_mem, Pointer.to(imgs.imgMem));
        CL.clSetKernelArg(kernel, 1, Sizeof.cl_mem, Pointer.to(outR.cl_data));
        CL.clSetKernelArg(kernel, 2, Sizeof.cl_mem, Pointer.to(outD.cl_data));
        CL.clSetKernelArg(kernel, 3, Sizeof.cl_float * imgs.w, null);
        CL.clSetKernelArg(kernel, 4, Sizeof.cl_float * imgs.w, null);
        CL.clSetKernelArg(kernel, 5, Sizeof.cl_int * params.length, Pointer.to(params));

        cl_event event = new cl_event();

        long localWorkSize[] = {imgs.w};
        long globalWorkSize[] = {localWorkSize[0] * imgs.numOfImages};
        cl_event[] waitForEventsArray = waitForEvents.toArray(new cl_event[waitForEvents.size()]);
        CL.clEnqueueNDRangeKernel(commandQueue, kernel, 1, null, globalWorkSize, localWorkSize, waitForEventsArray.length, waitForEventsArray, event);
        outR.cl_event = event;
        outD.cl_event = event;

//        waitForEvents.add(event);
//        print_cl_mem_Float(outR.cl_data, imgs.numOfImages, waitForEvents);
//        print_cl_mem_byte(imgs.imgMem, 10, waitForEvents);
        ArrayList<CL_Feature> features = new ArrayList<>();
        features.add(outR);
        features.add(outD);
        return features;
    }

    private ArrayList<CL_Feature> correlation(CL_PlaneSet imgs, ArrayList<cl_event> waitForEvents, CL_Feature mean, CL_Feature moment2) {
        cl_kernel kernel = clCorrelationKernel;
        CL_Feature outR = new CL_Feature();
        CL_Feature outD = new CL_Feature();

        outR.name = Feature.Names.CorrelationR;
        outD.name = Feature.Names.CorrelationD;

        outR.cl_data = CL.clCreateBuffer(context, CL.CL_MEM_READ_WRITE, imgs.numOfImages * Sizeof.cl_float, null, null);
        outD.cl_data = CL.clCreateBuffer(context, CL.CL_MEM_READ_WRITE, imgs.numOfImages * Sizeof.cl_float, null, null);

        int numelPixels = imgs.w * imgs.h;
        int params[] = {imgs.numOfImages, numelPixels, imgs.w, 0};
        CL.clSetKernelArg(kernel, 0, Sizeof.cl_mem, Pointer.to(imgs.imgMem));
        CL.clSetKernelArg(kernel, 1, Sizeof.cl_mem, Pointer.to(mean.cl_data));
        CL.clSetKernelArg(kernel, 2, Sizeof.cl_mem, Pointer.to(moment2.cl_data));
        CL.clSetKernelArg(kernel, 3, Sizeof.cl_mem, Pointer.to(outR.cl_data));
        CL.clSetKernelArg(kernel, 4, Sizeof.cl_mem, Pointer.to(outD.cl_data));
        CL.clSetKernelArg(kernel, 5, Sizeof.cl_float * imgs.w, null);
        CL.clSetKernelArg(kernel, 6, Sizeof.cl_float * imgs.w, null);
        CL.clSetKernelArg(kernel, 7, Sizeof.cl_int * params.length, Pointer.to(params));

        cl_event event = new cl_event();

        long localWorkSize[] = {imgs.w};
        long globalWorkSize[] = {localWorkSize[0] * imgs.numOfImages};
        cl_event[] waitForEventsArray = waitForEvents.toArray(new cl_event[waitForEvents.size()]);
        CL.clEnqueueNDRangeKernel(commandQueue, kernel, 1, null, globalWorkSize, localWorkSize, waitForEventsArray.length, waitForEventsArray, event);
        outR.cl_event = event;
        outD.cl_event = event;

        ArrayList<CL_Feature> features = new ArrayList<>();
        features.add(outR);
        features.add(outD);
        return features;
    }    
    
    private ArrayList<CL_Feature> kurtosisSkewness(CL_PlaneSet imgs, ArrayList<cl_event> waitForEvents, CL_Feature means) {
        cl_kernel kernel = clKurtosisSkewnessKernal;
        CL_Feature kurtosis = new CL_Feature();
        CL_Feature skewness = new CL_Feature();

        kurtosis.name = Feature.Names.Kurtosis;
        skewness.name = Feature.Names.Skewness;

        int numelPixels = imgs.w * imgs.h;
        int params[] = {imgs.numOfImages, numelPixels, imgs.w, 0};

        kurtosis.cl_data = CL.clCreateBuffer(context, CL.CL_MEM_READ_WRITE, imgs.numOfImages * Sizeof.cl_float, null, null);
        skewness.cl_data = CL.clCreateBuffer(context, CL.CL_MEM_READ_WRITE, imgs.numOfImages * Sizeof.cl_float, null, null);

        CL.clSetKernelArg(kernel, 0, Sizeof.cl_mem, Pointer.to(imgs.imgMem));
        CL.clSetKernelArg(kernel, 1, Sizeof.cl_mem, Pointer.to(means.cl_data));
        CL.clSetKernelArg(kernel, 2, Sizeof.cl_mem, Pointer.to(kurtosis.cl_data));
        CL.clSetKernelArg(kernel, 3, Sizeof.cl_mem, Pointer.to(skewness.cl_data));
        CL.clSetKernelArg(kernel, 4, Sizeof.cl_int4, Pointer.to(params));

        cl_event event = new cl_event();
        long localWorkSize[] = {imgs.w, 1};
        long globalWorkSize[] = {localWorkSize[0] * imgs.numOfImages, 1};
        cl_event[] waitForEventsArray = waitForEvents.toArray(new cl_event[waitForEvents.size()]);
        CL.clEnqueueNDRangeKernel(commandQueue, kernel, 1, null, globalWorkSize, localWorkSize, waitForEventsArray.length, waitForEventsArray, event);
        kurtosis.cl_event = event;
        skewness.cl_event = event;

        ArrayList<CL_Feature> features = new ArrayList<>();
        features.add(kurtosis);
        features.add(skewness);

        return features;
    }

    private CL_Feature magnitude(CL_PlaneSet imgs, ArrayList<cl_event> waitForEvents, CL_Feature f1, CL_Feature f2) {
        cl_kernel kernel = clMagnitudeKernel;
        CL_Feature fOut = new CL_Feature();
        fOut.cl_data = CL.clCreateBuffer(context, CL.CL_MEM_READ_WRITE, imgs.numOfImages * Sizeof.cl_float, null, null);

        int params[] = {imgs.numOfImages, 0};
        CL.clSetKernelArg(kernel, 0, Sizeof.cl_mem, Pointer.to(f1.cl_data));
        CL.clSetKernelArg(kernel, 1, Sizeof.cl_mem, Pointer.to(f2.cl_data));
        CL.clSetKernelArg(kernel, 2, Sizeof.cl_mem, Pointer.to(fOut.cl_data));
        CL.clSetKernelArg(kernel, 3, Sizeof.cl_int * params.length, Pointer.to(params));

        cl_event event = new cl_event();

        long localWorkSize[] = {16};
        long globalWorkSize[] = {round(localWorkSize[0], imgs.numOfImages)};
        cl_event[] waitForEventsArray = waitForEvents.toArray(new cl_event[waitForEvents.size()]);
        CL.clEnqueueNDRangeKernel(commandQueue, kernel, 1, null, globalWorkSize, localWorkSize, waitForEventsArray.length, waitForEventsArray, event);
        fOut.cl_event = event;
        return fOut;
    }    
    
    private CL_Feature fSum(CL_PlaneSet imgs, ArrayList<cl_event> waitForEvents, CL_Feature f1, CL_Feature f2) {
        cl_kernel kernel = clSumKernel;
        CL_Feature fOut = new CL_Feature();
        fOut.cl_data = CL.clCreateBuffer(context, CL.CL_MEM_READ_WRITE, imgs.numOfImages * Sizeof.cl_float, null, null);

        int params[] = {imgs.numOfImages, 0};
        CL.clSetKernelArg(kernel, 0, Sizeof.cl_mem, Pointer.to(f1.cl_data));
        CL.clSetKernelArg(kernel, 1, Sizeof.cl_mem, Pointer.to(f2.cl_data));
        CL.clSetKernelArg(kernel, 2, Sizeof.cl_mem, Pointer.to(fOut.cl_data));
        CL.clSetKernelArg(kernel, 3, Sizeof.cl_int * params.length, Pointer.to(params));

        cl_event event = new cl_event();

        long localWorkSize[] = {16};
        long globalWorkSize[] = {round(localWorkSize[0], imgs.numOfImages)};
        cl_event[] waitForEventsArray = waitForEvents.toArray(new cl_event[waitForEvents.size()]);
        CL.clEnqueueNDRangeKernel(commandQueue, kernel, 1, null, globalWorkSize, localWorkSize, waitForEventsArray.length, waitForEventsArray, event);
        fOut.cl_event = event;

        return fOut;
    }    
    
}
